#include <Dynamics/ExtForce/gravity_cuda.h>
#include <cuda_block_size_set.h>

namespace PhysLeo
{

template<typename T>
__global__ static void K_GravityCuda_addGravityForce(glm::tvec3<T>* force_ptr, T* mass_ptr, glm::tvec3<T> gravity_acceleration, int size)
{
    int pid = threadIdx.x + (blockIdx.x * blockDim.x);
    if (pid >= size) return;
    const auto vert_id = pid;

    force_ptr[vert_id] += mass_ptr[vert_id] * gravity_acceleration;
}
    
template<typename T>
void GravityCuda<T>::addGravityForce(std::shared_ptr<BufferData<glm::tvec3<T>>> force)
{
    auto force_ptr = force->ptrGpu();
    auto mass_ptr = GravityCpu<T>::ptr_mass_->ptrGpu();
    const int vert_num = GravityCpu<T>::ptr_mass_->size();

    K_GravityCuda_addGravityForce<T> << <numBlocks(vert_num), BLOCK_SIZE >> > (force_ptr, mass_ptr, GravityCpu<T>::gravity_acceleration_, vert_num);
    synchronCheck;
}

template class GravityCuda<float>;
template class GravityCuda<double>;

}